Repository navigation
TransferBench v1.71.00 - #359
Open
AtlantaPepsi wants to merge 24 commits into
Open
AtlantaPepsi wants to merge 24 commits into
AtlantaPepsi wants to merge 24 commits into
Conversation
Contributor
There was a problem hiding this comment.
Copilot review overview
🟡 Changes recommended
Unresolved critical and moderate correctness, portability, parsing, and build-target issues remain.
Get a fresh assessment by requesting another Copilot review.
Review effort: Lite
Findings: 4
Open (8)
What changed in this PR
Adds pingpong latency testing, NIC message-size validation, logical CU reporting fixes, and gfx1250-strict support.
Changes:
- Adds pingpong parsing, execution, timing, and latency presets.
- Adds NIC
max_msg_szvalidation and reporting. - Updates GPU architecture and TDM handling.
| File | Review summary |
|---|---|
src/header/TransferBench.hpp |
Critical and moderate issues in XCC handling, pingpong validation, NIC limits, subindex selection, dump parsing, and timing scaling. |
src/header/tdmCopy.h |
Critical architecture guard mismatch enables unsupported TDM targets. |
src/client/Utilities.hpp |
Reviewed result and topology utilities. |
src/client/Topology.hpp |
Reviewed NIC message-size reporting. |
src/client/Presets/Presets.hpp |
Reviewed latency preset registration. |
src/client/Presets/Latency.hpp |
Moderate issues with failed-run result handling and diagonal pair measurement. |
src/client/EnvVars.hpp |
Reviewed pingpong configuration variables. |
src/client/Client.cpp |
Reviewed pingpong transfer display updates. |
docs/install/build_from_source.rst |
Reviewed strict GPU target documentation. |
CMakeLists.txt |
Moderate issue: package build targets omit gfx1250-strict. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
| exeInfo.totalSubExecs += t.numSubExecs; | ||
| } else { | ||
| exeInfo.totalPingpong ++; | ||
| } | ||
| exeInfo.useSubIndices |= (t.exeSubIndex != -1 || (t.exeDevice.exeType == EXE_GPU_GFX && !cfg.gfx.prefXccTable.empty())); |
Contributor
There was a problem hiding this comment.
Copilot review overview
🟡 Changes recommended
Unresolved correctness issues remain in ping-pong execution, validation, parsing, and target gating.
Get a fresh assessment by requesting another Copilot review.
Review effort: Lite
Findings: 5
Open (9)
AtlantaPepsi
force-pushed
the
candidate-1.71
branch
from
September 23, 2026 21:06
e37e7de to
07e9f67
Compare
Contributor
There was a problem hiding this comment.
Copilot review overview
🟡 Changes recommended
Unresolved critical and moderate issues remain in pingpong execution, validation, result handling, and architecture gating.
Get a fresh assessment by requesting another Copilot review.
Review effort: Lite
Findings: 7
Open (12)
| #else | ||
| useSubIndexCount[exe]++; | ||
| int numSubIndices = GetNumExecutorSubIndices(exe); | ||
| if (subIndex >= numSubIndices) { |
Comment on lines
+6660
to
+6661
| dim3 const gridSize(xccDim, numPingpong, 1); | ||
| dim3 const blockSize(1); |
Comment on lines
3111
to
+3115
| if (t.numSubExecs <= 0) | ||
| errors.push_back({ERR_FATAL, "Transfer %d: # of subexecutors must be positive", i}); | ||
| else | ||
| else if (isPingpong) { | ||
| if (t.numSubExecs != 1) | ||
| errors.push_back({ERR_WARN, |
AtlantaPepsi
force-pushed
the
candidate-1.71
branch
from
September 28, 2026 15:27
07e9f67 to
91f99bd
Compare
Contributor
There was a problem hiding this comment.
Copilot review overview
🟡 Changes recommended
Unresolved correctness issues affect ping-pong execution, target guards, NIC validation, and package builds.
Review effort: Lite
Findings: 7
Open (11)
Comment on lines
+5562
to
+5567
| } else if (partnerMem.memIndex != exeDevice.exeIndex) { | ||
| if (System::Get().IsVerbose()) { | ||
| System::Get().Log("[INFO] Enabling pingpong peer access: GPU %d -> GPU %d\n", | ||
| exeDevice.exeIndex, partnerMem.memIndex); | ||
| } | ||
| ERR_CHECK(EnablePeerAccess(exeDevice.exeIndex, partnerMem.memIndex)); |
AtlantaPepsi
force-pushed
the
candidate-1.71
branch
from
October 8, 2026 13:47
91f99bd to
5ec784d
Compare
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
Unresolved critical and moderate issues remain in pingpong parsing and memory access, architecture guards, NIC validation, and latency failure propagation.
7 open findings
numSubExecsis never initialized because the parser only reads it in… · New
6 resolved since last review
🧠 Review effort: Lite
Give feedback about Copilot approvals in this survey to enter a drawing for a $150 gift card.
| p.srcMem[0] = nullptr; | ||
| p.srcMem[1] = nullptr; | ||
| p.localFlagMem = nullptr; | ||
| p.flagMem = static_cast<volatile uint8_t*>(static_cast<void*>(rss.dstMem[0])); |
| // Expand pong half | ||
| std::vector<Transfer> pongTransfers; | ||
| for (int r = 0; r < numRanks; r++) { | ||
| if (!RecursiveWildcardTransferExpansion(pongWct, r, numBytes, numSubExecs, pongTransfers)) |
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
Unresolved critical and moderate correctness issues remain in ping-pong synchronization, target guards, validation, and error propagation.
8 open findings
numSubExecsis never initialized because the parser only reads it in…
🧠 Review effort: Lite
Give feedback about Copilot approvals in this survey to enter a drawing for a $150 gift card.
Comment on lines
+6566
to
+6570
| // Advance one stride, plus an extra stride every hp laps so that a slot is never | ||
| // revisited an even number of laps later (which would leave a stale matching value) | ||
| off += sx; if (off >= y) off -= y; | ||
| if (--hopCnt == 0) { | ||
| hopCnt = hp; |
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
Unresolved critical and moderate correctness issues remain in pingpong execution, target handling, validation, and probe tooling.
10 open findings
numSubExecsis never initialized because the parser only reads it in…
🧠 Review effort: Lite
Give feedback about Copilot approvals in this survey to enter a drawing for a $150 gift card.
Comment on lines
6810
to
+6812
| for (int i = 0; i < exeInfo.resources.size(); i++) { | ||
| TransferResources& rss = exeInfo.resources[i]; | ||
| if (rss.numLaps != 0) continue; |
Comment on lines
+11
to
+30
| ## Single-node bandwidth presets | ||
|
|
||
| | Preset | Purpose | | ||
| |---|---| | ||
| | `a2a` | All-to-all parallel transfers between every pair of GPUs. | | ||
| | `a2asweep` | GFX-based a2a swept across CU counts and unroll factors (`MEM_TYPE`, `NUM_SUB_EXECS`). | | ||
| | `bmasweep` | Compares DMA vs. Batched-DMA for one-to-many copies (HIP 7.1 / CUDA 12.8+). | | ||
| | `gfxsweep` | Sweeps GFX kernel options for one Transfer. | | ||
| | `hbm` | Local HBM read bandwidth on each GPU. | | ||
| | `healthcheck` | Quick correctness/perf health check (AMD MI300 series only). | | ||
| | `one2all` | All subsets of parallel transfers from one GPU to all others. | | ||
| | `p2p` | Peer-to-peer device-memory matrix between every GPU pair. | | ||
| | `pcopy` | Parallel copies from a single GPU to other GPUs. | | ||
| | `rsweep` | Random sweep through Transfer combinations. | | ||
| | `rwrite` | Parallel remote writes from a single GPU to others. | | ||
| | `scaling` | Scaling test: one GPU → all others, varying SEs, mem types (`CPU_MEM_TYPE`, `GPU_MEM_TYPE`). | | ||
| | `schmoo` | Local/remote read/write/copy scaling between two GPUs. | | ||
| | `smoketest` | Quick DMA/GFX correctness sweep. | | ||
| | `sweep` | Ordered sweep through Transfer combinations. | | ||
| | `wallclock` | Compares wallclock counters across XCCs within one GPU. | |
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
Unresolved security-scan allowlist gaps and functional correctness issues remain.
12 open findings
numSubExecsis never initialized because the parser only reads it in…
Previously missed (3)
In code that hasn't changed since last review
.claude/skills/transferbench-debug/examples/topology-probe.sh:24
With set -o pipefail, truncating the tee pipeline with head can close the pipe while the binary is still writing; tee/the binary then exits on SIGPIPE and the script aborts before completing the probe. Use a reader that consumes the full stream (for example sed -n '1,60p') or capture first and truncate afterward.
.claude/skills/transferbench-run/references/presets.md:25
This reference presents these rows as runnable presets, but neither pcopy nor rwrite is registered in the current presetFuncMap (src/client/Presets/Presets.hpp:68-99), so following either command fails instead of launching a benchmark. Remove them or register the corresponding presets before documenting them here.
CMakeLists.txt:224
Adding the target only to CMake's default list does not add it to the package build path: build_packages_local.sh passes its separate DEFAULT_GPU_TARGETS list to -DGPU_TARGETS, and that list still omits both gfx1250 and gfx1250-strict. Packages built without an explicit GPU_TARGETS override therefore will not contain the target advertised by this change; update the packaging defaults or explicitly document the scope.
🧠 Review effort: Lite
Give feedback about Copilot approvals in this survey to enter a drawing for a $150 gift card.
Comment on lines
+35
to
+44
| [[allowlists]] | ||
| # Allowlist by location only where a real first-party secret structurally | ||
| # can't live. Dependency lock files qualify: their contents are generated | ||
| # from a manifest, and the high-entropy strings they carry are artifact | ||
| # digests rather than credentials. Add vendored third-party trees here as | ||
| # callers bring them in, one explicit path per entry. | ||
| description = "Generated dependency lock files" | ||
| paths = [ | ||
| '''.*\.lock$''', | ||
| ] |
Comment on lines
+53
to
+55
| '''(?i)\bINVALID_TOKEN\s*=''', | ||
| '''(?i)\bwrong_secret\s*=''', | ||
| '''(?i)\bsecret\s*=''' |
* fix (client): pre-resolve master address in LaunchTransferBench The host list was forwarded verbatim as TB_MASTER_ADDR, so an ssh_config alias that the local SSH client understands would fail getaddrinfo() on the workers, leaving rank 0 waiting on connections that never arrive. Expand the entry via ssh -G and prefer a literal IPv4, since workers resolve the master address remotely and only over AF_INET. * Potential fix for pull request finding Co-authored-by: Copilot Autofix powered by AI <[email protected]> * Apply suggestion from @nileshnegi * Potential fix for pull request finding Co-authored-by: Copilot Autofix powered by AI <[email protected]> --------- Co-authored-by: Copilot Autofix powered by AI <[email protected]>
Copies standardised security scanning config from ROCm/rocm-repo-template. Co-authored-by: haribabug <[email protected]>
AtlantaPepsi
force-pushed
the
candidate-1.71
branch
from
October 15, 2026 21:40
4193214 to
fa551ee
Compare
Contributor
There was a problem hiding this comment.
🔵 Needs a closer look
Unresolved correctness, documentation, packaging, and security-configuration issues remain.
28 open findings
numSubExecsis never initialized because the parser only reads it in…
And 8 more that still need to be addressed.
🧠 Review effort: Lite
AtlantaPepsi
force-pushed
the
candidate-1.71
branch
from
October 16, 2026 17:27
fa551ee to
614e3d3
Compare
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
Unresolved correctness, packaging, documentation, and security-configuration issues remain, including a critical ping-pong capacity concern.
30 open findings
numSubExecsis never initialized because the parser only reads it in…
And 10 more that still need to be addressed.
🧠 Review effort: Lite
Comment on lines
+3158
to
3164
| int totalTotalSubExecs = totalSubExecs[exeDevice] + totalPingpong[exeDevice]; | ||
| if (totalTotalSubExecs > numGpuSubExec) | ||
| errors.push_back({ERR_WARN, | ||
| "GPU %d requests %d total %s however only %d available. " | ||
| "Serialization will occur", | ||
| exeDevice.exeIndex, totalSubExecs[exeDevice], | ||
| exeDevice.exeIndex, totalTotalSubExecs, | ||
| cfg.gfx.seType == 0 ? "CUs" : "warps", numGpuSubExec}); |
| Plus two safe inspections of any preset/config: | ||
|
|
||
| ```bash | ||
| ./TransferBench dryrun "<expression>" # validate parsing, expand wildcards |
AtlantaPepsi
force-pushed
the
candidate-1.71
branch
from
October 16, 2026 19:09
614e3d3 to
2950ba8
Compare
Contributor
There was a problem hiding this comment.
🔵 Needs a closer look
Unresolved functional, packaging, CLI-example, and documentation issues remain.
30 open findings
numSubExecsis never initialized because the parser only reads it in…
And 10 more that still need to be addressed.
🧠 Review effort: Lite
AtlantaPepsi
force-pushed
the
candidate-1.71
branch
from
October 16, 2026 20:36
2950ba8 to
53c3358
Compare
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
Unresolved critical and moderate findings remain in target gating, pingpong execution, packaging, launcher behavior, and documentation/examples.
32 open findings
numSubExecsis never initialized because the parser only reads it in…
And 12 more that still need to be addressed.
🧠 Review effort: Lite
| @@ -788,7 +857,7 @@ __device__ __forceinline__ uint32_t GetXccId() | |||
| uint32_t xccReg = 0; | |||
| asm volatile("s_getreg_b32 %0, hwreg(HW_REG_XCC_ID)" : "=s"(xccReg)); | |||
| return xccReg & 0xF; | |||
| #elif defined(__gfx1250__) | |||
| #elif defined(__GFX12__) | |||
Comment on lines
+101
to
+103
| 1. `dryrun "<expression>"` first — most parser-level bugs surface here. | ||
| 2. Read the banner from `./TransferBench` with no args — confirms which memory types and executors are compiled in for this build. | ||
| 3. For DMA crashes: confirm exactly one SRC per Transfer. |
AtlantaPepsi
force-pushed
the
candidate-1.71
branch
from
October 16, 2026 21:24
53c3358 to
aaa114c
Compare
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
Fifteen final review comments remain unresolved, including moderate correctness, documentation, packaging, and security issues.
33 open issues
numSubExecsis never initialized because the parser only reads it in…
And 13 more that still need to be addressed.
Review effort: Lite
|
|
||
| # Default GPU targets baked into every package, regardless of GPU_FAMILY tarball. | ||
| DEFAULT_GPU_TARGETS="gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1200;gfx1201" | ||
| DEFAULT_GPU_TARGETS="gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201" |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
Technical Details
Test Plan
HSA_DISABLE_GFX12_STRICT=0 ./TransferBenchoutput of gfx1250-strict targetSHOW_ITERATIONSoutput CU ID againstCU_MASKbitsp2p_latencypreset: in/cross-domain, and AMD and Nvidia platformTest Result
p2p_latency example output
Concurrent transfer + pingpong hybrid run example output
Submission Checklist